// Enter your kernel in this window

typedef struct _BUFFER_INFO
{
	unsigned int Offset;
	unsigned int Size;
} BUFFER_INFO, *PBUFFER_INFO;

typedef struct _STREAM_BUFFER_INFO
{
	unsigned int OriginalSize;
	char Finished;
} STREAM_BUFFER_INFO, *PSTREAM_BUFFER_INFO;


#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
#pragma OPENCL EXTENSION cl_amd_printf : enable

/* The basic MD5 functions */
#define F(x, y, z)	((z) ^ ((x) & ((y) ^ (z))))
#define G(x, y, z)	((y) ^ ((z) & ((x) ^ (y))))
#define H(x, y, z)	((x) ^ (y) ^ (z))
#define I(x, y, z)	((y) ^ ((x) | ~(z)))

/* The MD5 transformation for all four rounds. */
#define STEP(f, a, b, c, d, x, t, s) \
(a) += f((b), (c), (d)) + (x) + (t); \
(a) = (((a) << (s)) | (((a) & 0xffffffff) >> (32 - (s)))); \
    (a) += (b);

void md5_round(const unsigned int buffer[64], unsigned int* internal_state/*, const unsigned int funcInfoLen, char* funcInfo*/)
{
	unsigned int a, b, c, d;
	a = internal_state[0];
	b = internal_state[1];
	c = internal_state[2];
	d = internal_state[3];

	/* Round 1 */
	STEP(F, a, b, c, d, (buffer[(0)]), 0xd76aa478, 7)
	STEP(F, d, a, b, c, (buffer[(1)]), 0xe8c7b756, 12)
	STEP(F, c, d, a, b, (buffer[(2)]), 0x242070db, 17)
	STEP(F, b, c, d, a, (buffer[(3)]), 0xc1bdceee, 22)
	STEP(F, a, b, c, d, (buffer[(4)]), 0xf57c0faf, 7)
	STEP(F, d, a, b, c, (buffer[(5)]), 0x4787c62a, 12)
	STEP(F, c, d, a, b, (buffer[(6)]), 0xa8304613, 17)
	STEP(F, b, c, d, a, (buffer[(7)]), 0xfd469501, 22)
	STEP(F, a, b, c, d, (buffer[(8)]), 0x698098d8, 7)
	STEP(F, d, a, b, c, (buffer[(9)]), 0x8b44f7af, 12)
	STEP(F, c, d, a, b, (buffer[(10)]), 0xffff5bb1, 17)
	STEP(F, b, c, d, a, (buffer[(11)]), 0x895cd7be, 22)
	STEP(F, a, b, c, d, (buffer[(12)]), 0x6b901122, 7)
	STEP(F, d, a, b, c, (buffer[(13)]), 0xfd987193, 12)
	STEP(F, c, d, a, b, (buffer[(14)]), 0xa679438e, 17)
	STEP(F, b, c, d, a, (buffer[(15)]), 0x49b40821, 22)

	/* Round 2 */
	STEP(G, a, b, c, d, (buffer[(1)]), 0xf61e2562, 5)
	STEP(G, d, a, b, c, (buffer[(6)]), 0xc040b340, 9)
	STEP(G, c, d, a, b, (buffer[(11)]), 0x265e5a51, 14)
	STEP(G, b, c, d, a, (buffer[(0)]), 0xe9b6c7aa, 20)
	STEP(G, a, b, c, d, (buffer[(5)]), 0xd62f105d, 5)
	STEP(G, d, a, b, c, (buffer[(10)]), 0x02441453, 9)
	STEP(G, c, d, a, b, (buffer[(15)]), 0xd8a1e681, 14)
	STEP(G, b, c, d, a, (buffer[(4)]), 0xe7d3fbc8, 20)
	STEP(G, a, b, c, d, (buffer[(9)]), 0x21e1cde6, 5)
	STEP(G, d, a, b, c, (buffer[(14)]), 0xc33707d6, 9)
	STEP(G, c, d, a, b, (buffer[(3)]), 0xf4d50d87, 14)
	STEP(G, b, c, d, a, (buffer[(8)]), 0x455a14ed, 20)
	STEP(G, a, b, c, d, (buffer[(13)]), 0xa9e3e905, 5)
	STEP(G, d, a, b, c, (buffer[(2)]), 0xfcefa3f8, 9)
	STEP(G, c, d, a, b, (buffer[(7)]), 0x676f02d9, 14)
	STEP(G, b, c, d, a, (buffer[(12)]), 0x8d2a4c8a, 20)

	/* Round 3 */
	STEP(H, a, b, c, d, (buffer[(5)]), 0xfffa3942, 4)
	STEP(H, d, a, b, c, (buffer[(8)]), 0x8771f681, 11)
	STEP(H, c, d, a, b, (buffer[(11)]), 0x6d9d6122, 16)
	STEP(H, b, c, d, a, (buffer[(14)]), 0xfde5380c, 23)
	STEP(H, a, b, c, d, (buffer[(1)]), 0xa4beea44, 4)
	STEP(H, d, a, b, c, (buffer[(4)]), 0x4bdecfa9, 11)
	STEP(H, c, d, a, b, (buffer[(7)]), 0xf6bb4b60, 16)
	STEP(H, b, c, d, a, (buffer[(10)]), 0xbebfbc70, 23)
	STEP(H, a, b, c, d, (buffer[(13)]), 0x289b7ec6, 4)
	STEP(H, d, a, b, c, (buffer[(0)]), 0xeaa127fa, 11)
	STEP(H, c, d, a, b, (buffer[(3)]), 0xd4ef3085, 16)
	STEP(H, b, c, d, a, (buffer[(6)]), 0x04881d05, 23)
	STEP(H, a, b, c, d, (buffer[(9)]), 0xd9d4d039, 4)
	STEP(H, d, a, b, c, (buffer[(12)]), 0xe6db99e5, 11)
	STEP(H, c, d, a, b, (buffer[(15)]), 0x1fa27cf8, 16)
	STEP(H, b, c, d, a, (buffer[(2)]), 0xc4ac5665, 23)

	/* Round 4 */
	STEP(I, a, b, c, d, (buffer[(0)]), 0xf4292244, 6)
	STEP(I, d, a, b, c, (buffer[(7)]), 0x432aff97, 10)
	STEP(I, c, d, a, b, (buffer[(14)]), 0xab9423a7, 15)
	STEP(I, b, c, d, a, (buffer[(5)]), 0xfc93a039, 21)
	STEP(I, a, b, c, d, (buffer[(12)]), 0x655b59c3, 6)
	STEP(I, d, a, b, c, (buffer[(3)]), 0x8f0ccc92, 10)
	STEP(I, c, d, a, b, (buffer[(10)]), 0xffeff47d, 15)
	STEP(I, b, c, d, a, (buffer[(1)]), 0x85845dd1, 21)
	STEP(I, a, b, c, d, (buffer[(8)]), 0x6fa87e4f, 6)
	STEP(I, d, a, b, c, (buffer[(15)]), 0xfe2ce6e0, 10)
	STEP(I, c, d, a, b, (buffer[(6)]), 0xa3014314, 15)
	STEP(I, b, c, d, a, (buffer[(13)]), 0x4e0811a1, 21)
	STEP(I, a, b, c, d, (buffer[(4)]), 0xf7537e82, 6)
	STEP(I, d, a, b, c, (buffer[(11)]), 0xbd3af235, 10)
	STEP(I, c, d, a, b, (buffer[(2)]), 0x2ad7d2bb, 15)
	STEP(I, b, c, d, a, (buffer[(9)]), 0xeb86d391, 21)

	internal_state[0] = a + internal_state[0];
	internal_state[1] = b + internal_state[1];
	internal_state[2] = c + internal_state[2];
	internal_state[3] = d + internal_state[3];
}

void md5(__global const char* dataBuffer, unsigned int size, unsigned int* restrict result)
{
	unsigned int i;
	unsigned int bytes_left;
	char buffer[64];
	
	result[0] = 0x67452301;
	result[1] = 0xefcdab89;
	result[2] = 0x98badcfe;
	result[3] = 0x10325476;

	for (bytes_left = size; bytes_left >= 64;
		bytes_left -= 64, dataBuffer = &dataBuffer[64])
	{
		for (i = 0; i < 64; ++i)
		{
			buffer[i] = dataBuffer[i];
		}
		/*printf("md5_round before (end of func):\n\tbuffer:");
		for (i=0; i<64; i++)
		{
			printf("%x", buffer[i]);
		}
		printf("\n\tresult: %x%x%x%x", result[0], result[1], result[2], result[3]);*/
		md5_round((const unsigned int*)buffer, result);
		//printf("md5_after before (end of func):\n\tbuffer:");
		/*for (i=0; i<64; i++)
		{
			printf("%x", buffer[i]);
		}
		printf("\n\tresult: %x%x%x%x", result[0], result[1], result[2], result[3]);*/
	}

	for (i = 0; i < bytes_left; i++)
	{
		buffer[i] = dataBuffer[i];
	}

	buffer[bytes_left++] = 0x80;

	if (bytes_left <= 56)
	{
		for (i = bytes_left; i < 56; buffer[i++] = 0);
	}
	else
	{
		// If we have to pad enough to roll past this round.
		for (i = bytes_left; i < 64; buffer[i++] = 0);
		/*printf("md5_round before (end of func):\n\tbuffer:");
		for (i=0; i<64; i++)
		{
			printf("%x", buffer[i]);
		}
		printf("\n\tresult: %x%x%x%x", result[0], result[1], result[2], result[3]);*/
		md5_round((const unsigned int*)buffer, result);
		/*printf("md5_round after (end of func):\n\tbuffer:");
		for (i=0; i<64; i++)
		{
			printf("%x", buffer[i]);
		}
		printf("\n\tresult: %x%x%x%x", result[0], result[1], result[2], result[3]);*/
		for (i = 0; i < 56; buffer[i++] = 0);
	}

	unsigned long* len_ptr = (unsigned long*)&buffer[56];
	*len_ptr = size * 8;
	/*printf("md5_round before (end of func):\n\tbuffer:");
	for (i=0; i<64; i++)
	{
		printf("%x", buffer[i]);
	}
	printf("\n\tresult: %x%x%x%x", result[0], result[1], result[2], result[3]);*/
	md5_round((const unsigned int*)buffer, result);
	/**printf("md5_round after (end of func):\n\tbuffer:");
	for (i=0; i<64; i++)
	{
		printf("%x", buffer[i]);
	}
	printf("\n\tresult: %x%x%x%x", result[0], result[1], result[2], result[3]);*/
}

__kernel void calc_md5_items(__global const char* dataBuffer, __global const BUFFER_INFO* info, __global unsigned int* hashBuffer)
{
	unsigned int bufferIndex = get_global_id(0);
	unsigned int hashResult[4];

	dataBuffer += info[bufferIndex].Offset;

	md5(dataBuffer, info[bufferIndex].Size, hashResult);

	for (unsigned int i = 0; i < 4; ++i)
	{
		__global unsigned int* hash = &hashBuffer[bufferIndex * 4];
		hash[i] = hashResult[i];
	}
}

bool md5_start(__global const char* dataBuffer, unsigned int size, unsigned int* restrict result)
{
	unsigned int i;
	unsigned int bytes_left;
	unsigned int current_iteration = 1024;
	char buffer[64];
	
	result[0] = 0x67452301;
	result[1] = 0xefcdab89;
	result[2] = 0x98badcfe;
	result[3] = 0x10325476;

	//printf("\n\nmd6_start function started\n");
	
	for (bytes_left = size; (current_iteration > 0 && (bytes_left >= 64));
		bytes_left -= 64, dataBuffer = &dataBuffer[64], --current_iteration)
	{
		for (i = 0; i < 64; ++i)
		{
			buffer[i] = dataBuffer[i];
		}
		/*
		printf("md5_round before:\nbuffer:");
		for (i=0; i<64; i++)
		{
			printf("%x", buffer[i]);
		}
		printf("\nresult: %x%x%x%x\n", result[0], result[1], result[2], result[3]);*/
		md5_round((const unsigned int*)buffer, result);
		/*printf("md5_round after:\nbuffer:");
		for (i=0; i<64; i++)
		{
			printf("%x", buffer[i]);
		}
		printf("\nresult: %x%x%x%x\n", result[0], result[1], result[2], result[3]);*/
	}

	if ((current_iteration == 0) && (bytes_left > 0))
	{
		//printf("\nreturn false\n");
		return (false);
	}

	for (i = 0; i < bytes_left; i++)
	{
		buffer[i] = dataBuffer[i];
	}

	buffer[bytes_left++] = 0x80;

	if (bytes_left <= 56)
	{
		for (i = bytes_left; i < 56; buffer[i++] = 0);
	}
	else
	{
		// If we have to pad enough to roll past this round.
		for (i = bytes_left; i < 64; buffer[i++] = 0);
		md5_round((const unsigned int*)buffer, result);
		for (i = 0; i < 56; buffer[i++] = 0);
	}

	unsigned long* len_ptr = (unsigned long*)&buffer[56];
	*len_ptr = size * 8;
	md5_round((const unsigned int*)buffer, result);

	//printf("\nmd5_start function finished\n\n");
	return (true);
}

bool md5_more(__global const char* dataBuffer, unsigned int size, unsigned int originalSize, unsigned int* restrict result)
{
	unsigned int i;
	unsigned int bytes_left = size;
	// todo: change magic number
	unsigned int current_iteration = 1024;
	char buffer[64];
	/*
	printf("\n\nmd5_more function started\n");

	printf("md5_more start of function:\n");
	printf("size: %d\n", size);
	printf("\n\tbuffer:\n");
	for (i=0; i<size; i++)
	{
		printf("%x", dataBuffer[i]);
	}
	printf("\n\tresult: %x%x%x%x\n", result[0], result[1], result[2], result[3]);
	*/
	if (size >= 64)
	{
		//printf("\nsize >= 64\n");
		for (bytes_left = size; ((current_iteration > 0) && (bytes_left >= 64));
			bytes_left -= 64, dataBuffer = &dataBuffer[64], --current_iteration)
		{
			for (i = 0; i < 64; ++i)
			{
				buffer[i] = dataBuffer[i];
			}
			
			/*printf("md5_round before:\n\tbuffer:");
			for (i=0; i<64; i++)
			{
				printf("%x", buffer[i]);
			}*/
			//printf("\n\tresult: %x%x%x%x", result[0], result[1], result[2], result[3]);
			md5_round((const unsigned int*)buffer, result);
			/*printf("\nmd5_round after:\n\tbuffer:");
			for (i=0; i<64; i++)
			{
				printf("%x", buffer[i]);
			}
			printf("\nresult: %x%x%x%x\n", result[0], result[1], result[2], result[3]);*/
		}

		if ((current_iteration == 0) && (bytes_left > 0))
		{
			//printf("\nreturn false\n");
			return (false);
		}
	}
 
	for (i = 0; i < bytes_left; i++)
	{
		buffer[i] = dataBuffer[i];
	}

	buffer[bytes_left++] = 0x80;

	if (bytes_left <= 56)
	{
		for (i = bytes_left; i < 56; buffer[i++] = 0);
	}
	else
	{
		// If we have to pad enough to roll past this round.
		for (i = bytes_left; i < 64; buffer[i++] = 0);
		md5_round((const unsigned int*)buffer, result);
		for (i = 0; i < 56; buffer[i++] = 0);
	}

	unsigned long* len_ptr = (unsigned long*)&buffer[56];
	*len_ptr = originalSize * 8;
	/*//printf("md5_round before (end of func):\n\tbuffer:");
	for (i=0; i<64; i++)
	{
		printf("%x", buffer[i]);
	}*/
	//printf("\n\tresult: %x%x%x%x", result[0], result[1], result[2], result[3]);
	md5_round((const unsigned int*)buffer, result);
	//printf("md5_round after (end of func):\n\tbuffer:");
	/*for (i=0; i<64; i++)
	{
		printf("%x", buffer[i]);
	}*/
	//printf("\n\tresult: %x%x%x%x", result[0], result[1], result[2], result[3]);

	//printf("\n\nmd5_more function finished\n");
	return (true);
}

__kernel void calc_md5_stream_start(__global const char* dataBuffer, __global const BUFFER_INFO* bufferInfo, __global STREAM_BUFFER_INFO* streamInfo, __global unsigned int* hashBuffer)
{
	unsigned int bufferIndex = get_global_id(0);
	unsigned int hashResult[4];
	__global unsigned int* hash = &hashBuffer[bufferIndex * 4];
	bool done;

	bufferInfo = &bufferInfo[bufferIndex];
	streamInfo = &streamInfo[bufferIndex];
	dataBuffer += bufferInfo->Offset;

	done = md5_start(dataBuffer, bufferInfo->Size, hashResult);

	if (done)
	{
		streamInfo->Finished = 1;
	}

	for (unsigned int i = 0; i < 4; ++i)
	{
		hash[i] = hashResult[i];
	}
}

__kernel void calc_md5_stream_more(__global const char* dataBuffer, __global const BUFFER_INFO* bufferInfo, __global STREAM_BUFFER_INFO* streamInfo, __global unsigned int* hashBuffer)
{
	unsigned int bufferIndex = get_global_id(0);
	unsigned int hashResult[4];
	__global unsigned int* hash = &hashBuffer[bufferIndex * 4];
	bool done;

	streamInfo = &streamInfo[bufferIndex];

	if (streamInfo->Finished == 0)
	{
		bufferInfo = &bufferInfo[bufferIndex];
		dataBuffer += bufferInfo->Offset;

		for (unsigned int i = 0; i < 4; ++i)
		{
			hashResult[i] = hash[i];
		}

		done = md5_more(dataBuffer, bufferInfo->Size, streamInfo->OriginalSize, hashResult);

		if (done)
		{
			streamInfo->Finished = 1;
		}

		for (unsigned int i = 0; i < 4; ++i)
		{
			hash[i] = hashResult[i];
		}
	}
}